/******************************************************************************
** This file is part of the cl-jpeg project.
**
** Copyright (C) 2011-2013 Vladimir Antonenko
**
** This program is free software; you can redistribute it and/or modify it
** under the terms of the GNU General Public License as published by the
** Free Software Foundation; either version 2 of the License,
** or (at your option) any later version.
**
** This program is distributed in the hope that it will be useful,
** but WITHOUT ANY WARRANTY; without even the implied warranty of
** MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
** See the GNU General Public License for more details.
**
** You should have received a copy of the GNU General Public License along
** with this program; if not, write to the Free Software Foundation, Inc.
******************************************************************************/

typedef struct
{
	uchar b;
	uchar g;
	uchar r;
}
BGR;

__kernel void rgb2yuv(__global const BGR *data,
					  __global char *Y,
					  __global char *Cb,
					  __global char *Cr)
{
	unsigned id = get_global_id(0);
	unsigned i = id;
	uchar    R, G, B;
	char     y, cb, cr;

	B = data[i].b;
	G = data[i].g;
	R = data[i].r;
	
	y =          0.299f  * R + 0.587f  * G + 0.114f  * B;
	cb = 128.f - 0.1687f * R - 0.3313f * G + 0.5f    * B;
	cr = 128.f + 0.5f    * R - 0.4187f * G - 0.0813f * B;

	Y[i]  = y;
	Cb[i] = cb;
	Cr[i] = cr;
}

// integer vectored kernel
__kernel void i_rgb2yuv(__global const BGR *d,
					  __global char4 *Y,
					  __global char4 *Cb,
					  __global char4 *Cr)
{
	size_t id = get_global_id(0);
	size_t i = id*4;

	int4 B = (int4) (d[i+0].b, d[i+1].b, d[i+2].b, d[i+3].b);
	int4 G = (int4) (d[i+0].g, d[i+1].g, d[i+2].g, d[i+3].g);
	int4 R = (int4) (d[i+0].r, d[i+1].r, d[i+2].r, d[i+3].r);
	
	Y [id] = convert_char4((  32768 + 19595*R + 38470*G +  7471*B) >> 16);
	Cb[id] = convert_char4((8421376 - 11058*R - 21709*G + 32767*B) >> 16);
	Cr[id] = convert_char4((8421376 + 32767*R - 27438*G -  5329*B) >> 16);
}

// vectored kernel
__kernel void v_rgb2yuv
	(
		__global const uint   *BGR,
		__global       uchar4 *Y,
		__global       uchar4 *Cb,
		__global       uchar4 *Cr
	)
{
	// BGR 4x3 -> 3x4 shuffle positions
	const uchar16 sf = (uchar16)(0,3,6,9, 1,4,7,10, 2,5,8,11, 0,0,0,0);
	size_t        id = get_global_id(0);	// for 4 pixels
	size_t        i = id*3;					// 4 pixels occupy 3 ints

	uint4 t = (uint4)(vload3(i,BGR), 0);
	t = as_uint4(shuffle(as_uchar16(t), sf));

	float4 R = convert_float4(as_uchar4(t.s0));
	float4 G = convert_float4(as_uchar4(t.s1));
	float4 B = convert_float4(as_uchar4(t.s2));

	// write 4 pixels
	Y [id] = convert_uchar4(  0.5f + 0.299f  * R + 0.587f  * G + 0.114f  * B);
	Cb[id] = convert_uchar4(128.5f - 0.1687f * R - 0.3313f * G + 0.499f  * B);
	Cr[id] = convert_uchar4(128.5f + 0.499f  * R - 0.4187f * G - 0.0813f * B);
}

__kernel void v2_rgb2yuv
	(
		__global const BGR    *d,
		__global       uchar4 *yuv
	)
{
	size_t  id = get_global_id(0);
	size_t  i = id*4;
	float4  R, G, B;
	float4  y, cb, cr;

	R = (float4) (d[i+0].r, d[i+1].r, d[i+2].r, d[i+3].r);
	G = (float4) (d[i+0].g, d[i+1].g, d[i+2].g, d[i+3].g);
	B = (float4) (d[i+0].b, d[i+1].b, d[i+2].b, d[i+3].b);

	y =    0.5f + 0.299f  * R + 0.587f  * G + 0.114f  * B;
	cb = 128.5f - 0.1687f * R - 0.3313f * G + 0.499f  * B;
	cr = 128.5f + 0.499f  * R - 0.4187f * G - 0.0813f * B;

	yuv[id] = convert_uchar4( y);
	yuv[id+SIZEY*SIZEX/4] = convert_uchar4(cb);
	yuv[id+SIZEY*SIZEX/2] = convert_uchar4(cr);
}

// 8x8 output order
__kernel void v3_rgb2yuv
	(
		__global const BGR    *d,
		__global       uchar8 *yuv
	)
{
	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t n = y / 8;
	size_t i = y*get_global_size(0)*8 + x*8;
	size_t o = n*get_global_size(0)*8 + x*8 + (y%8);
	float8 R, G, B;
	float8 c1, c2, c3;

	R = (float8) (d[i+0].r, d[i+1].r, d[i+2].r, d[i+3].r, d[i+4].r, d[i+5].r, d[i+6].r, d[i+7].r);
	G = (float8) (d[i+0].g, d[i+1].g, d[i+2].g, d[i+3].g, d[i+4].g, d[i+5].g, d[i+6].g, d[i+7].g);
	B = (float8) (d[i+0].b, d[i+1].b, d[i+2].b, d[i+3].b, d[i+4].b, d[i+5].b, d[i+6].b, d[i+7].b);

	c1 =   0.5f + 0.299f  * R + 0.587f  * G + 0.114f  * B;
	c2 = 128.5f - 0.1687f * R - 0.3313f * G + 0.499f  * B;
	c3 = 128.5f + 0.499f  * R - 0.4187f * G - 0.0813f * B;

	yuv[o] = convert_uchar8(c1);
	yuv[o+SIZEY*SIZEX/8] = convert_uchar8(c2);
	yuv[o+SIZEY*SIZEX/4] = convert_uchar8(c3);
}

#if 1
inline uchar8 RGB2Y2(const ushort8 R0, const ushort8 G0, const ushort8 B0)
{
	const uint8 R = convert_uint8(R0);
	const uint8 G = convert_uint8(G0);
	const uint8 B = convert_uint8(B0);
	return convert_uchar8((  32768 + 19595*R + 38470*G +  7471*B) >> 16);
}
inline uchar8 RGB2Cb2(const short8 R0, const short8 G0, const short8 B0)
{
	const int8 R = convert_int8(R0);
	const int8 G = convert_int8(G0);
	const int8 B = convert_int8(B0);
	return convert_uchar8((8421376*4 - 11058*R - 21709*G + 32767*B) >> 18);
}
inline uchar8 RGB2Cr2(const short8 R0, const short8 G0, const short8 B0)
{
	const int8 R = convert_int8(R0);
	const int8 G = convert_int8(G0);
	const int8 B = convert_int8(B0);
	return convert_uchar8((8421376*4 + 32767*R - 27438*G -  5329*B) >> 18);
}
#else
inline uchar8 RGB2Y2(const ushort8 R, const ushort8 G, const ushort8 B)
{
	const ushort8 k0 = (ushort8) 0;
	const ushort8 kR = (ushort8) 19595;
	const ushort8 kG = (ushort8) 38470;
	const ushort8 kB = (ushort8) 7471;
	return convert_uchar8(mad_hi(kB,B, mad_hi(kG,G, mad_hi(kR,R, k0))));
}
inline uchar8 RGB2Cb2(const short8 R, const short8 G, const short8 B)
{
	const short8 k0 = (short8) (128<<2);
	const short8 kR = (short8) -11058;
	const short8 kG = (short8) -21709;
	const short8 kB = (short8)  32767;
	return convert_uchar8(mad_hi(kB,B, mad_hi(kG,G, mad_hi(kR,R, k0))) >> 2);
}
inline uchar8 RGB2Cr2(const short8 R, const short8 G, const short8 B)
{
	const short8 k0 = (short8) (128<<2);
	const short8 kR = (short8)  32767;
	const short8 kG = (short8) -27438;
	const short8 kB = (short8)  -5329;
	return convert_uchar8(mad_hi(kB,B, mad_hi(kG,G, mad_hi(kR,R, k0))) >> 2);
}
#endif

inline ushort4 hsum2(const ushort8 t)
{
	return (ushort4) (t.s0+t.s1, t.s2+t.s3, t.s4+t.s5, t.s6+t.s7);
}

// BGR to YUV conversion + chroma subsampling 4:2:0
// Input: 2D, dimension [SIZEX/16,SIZEY/2]
__kernel void v4_rgb2yuv420
	(
		__global const uint   *bgr,
		__global       uchar8 *yuv
	)
{
	const uchar16 mask = (uchar16)(2,5,8,11, 1,4,7,10, 0,3,6,9, 0,0,0,0);
	const size_t  x = get_global_id(0);
	size_t        y = get_global_id(1)*2;
	size_t        i = (y*get_global_size(0) + x)*4;
	size_t        o = ((y/8)*get_global_size(0) + x)*16 + (y%8);
	uint2         rgb[3][2][2];

	{
		uint4 t = (uint4)(vload3(i+0,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][0][0].lo = t.s0;
		rgb[1][0][0].lo = t.s1;
		rgb[2][0][0].lo = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+1,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][0][0].hi = t.s0;
		rgb[1][0][0].hi = t.s1;
		rgb[2][0][0].hi = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+2,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][0][1].lo = t.s0;
		rgb[1][0][1].lo = t.s1;
		rgb[2][0][1].lo = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+3,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][0][1].hi = t.s0;
		rgb[1][0][1].hi = t.s1;
		rgb[2][0][1].hi = t.s2;
	}
	// second line
	y += 1;
	i = (y*get_global_size(0) + x)*4;
	{
		uint4 t = (uint4)(vload3(i+0,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][1][0].lo = t.s0;
		rgb[1][1][0].lo = t.s1;
		rgb[2][1][0].lo = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+1,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][1][0].hi = t.s0;
		rgb[1][1][0].hi = t.s1;
		rgb[2][1][0].hi = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+2,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][1][1].lo = t.s0;
		rgb[1][1][1].lo = t.s1;
		rgb[2][1][1].lo = t.s2;
	}{
		uint4 t = (uint4)(vload3(i+3,bgr), 0);
		t = as_uint4(shuffle(as_uchar16(t), mask));
		rgb[0][1][1].hi = t.s0;
		rgb[1][1][1].hi = t.s1;
		rgb[2][1][1].hi = t.s2;
	}
	{
		ushort8 R = convert_ushort8(as_uchar8(rgb[0][0][0]));
		rgb[0][0][0] = as_uint2(hsum2(R));
		ushort8 G = convert_ushort8(as_uchar8(rgb[1][0][0]));
		rgb[1][0][0] = as_uint2(hsum2(G));
		ushort8 B = convert_ushort8(as_uchar8(rgb[2][0][0]));
		rgb[2][0][0] = as_uint2(hsum2(B));
		yuv[o] = RGB2Y2(R,G,B);
	}{
		ushort8 R = convert_ushort8(as_uchar8(rgb[0][0][1]));
		rgb[0][0][1] = as_uint2(hsum2(R));
		ushort8 G = convert_ushort8(as_uchar8(rgb[1][0][1]));
		rgb[1][0][1] = as_uint2(hsum2(G));
		ushort8 B = convert_ushort8(as_uchar8(rgb[2][0][1]));
		rgb[2][0][1] = as_uint2(hsum2(B));
		yuv[o+8] = RGB2Y2(R,G,B);
	}
	o = ((y/8)*get_global_size(0) + x)*16 + (y%8);
	{
		ushort8 R = convert_ushort8(as_uchar8(rgb[0][1][0]));
		rgb[0][1][0] = as_uint2(hsum2(R));
		ushort8 G = convert_ushort8(as_uchar8(rgb[1][1][0]));
		rgb[1][1][0] = as_uint2(hsum2(G));
		ushort8 B = convert_ushort8(as_uchar8(rgb[2][1][0]));
		rgb[2][1][0] = as_uint2(hsum2(B));
		yuv[o] = RGB2Y2(R,G,B);
	}{
		ushort8 R = convert_ushort8(as_uchar8(rgb[0][1][1]));
		rgb[0][1][1] = as_uint2(hsum2(R));
		ushort8 G = convert_ushort8(as_uchar8(rgb[1][1][1]));
		rgb[1][1][1] = as_uint2(hsum2(G));
		ushort8 B = convert_ushort8(as_uchar8(rgb[2][1][1]));
		rgb[2][1][1] = as_uint2(hsum2(B));
		yuv[o+8] = RGB2Y2(R,G,B);
	}

	// chroma
	{
		const size_t offset_cb = get_global_size(1)*get_global_size(0)*4;
		const size_t offset_cr = offset_cb + get_global_size(1)*get_global_size(0);

		y = get_global_id(1);
		o = ((y/8)*get_global_size(0) + x)*8 + (y%8);

		// RGB addition of the pair in vetical direction
		short8 R = ((short8)(as_short4(rgb[0][0][0]),as_short4(rgb[0][0][1])))
		         + ((short8)(as_short4(rgb[0][1][0]),as_short4(rgb[0][1][1])));

		short8 G = ((short8)(as_short4(rgb[1][0][0]),as_short4(rgb[1][0][1])))
		         + ((short8)(as_short4(rgb[1][1][0]),as_short4(rgb[1][1][1])));

		short8 B = ((short8)(as_short4(rgb[2][0][0]),as_short4(rgb[2][0][1])))
		         + ((short8)(as_short4(rgb[2][1][0]),as_short4(rgb[2][1][1])));

		yuv[offset_cb+o] = RGB2Cb2(R,G,B);
		yuv[offset_cr+o] = RGB2Cr2(R,G,B);
	}
}

// 2D chrome subsampling kernel
__kernel void subsample(
					__global const uchar8 *src1,
					__global const uchar8 *src2,
					__global uchar4 *dst1,
					__global uchar4 *dst2)
{
	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t i1 = y*2*SIZEX/8 + x;	// first line
	size_t i2 = i1 + SIZEX/8;		// secons line

	uchar4 avr1;
	uchar4 avr2;

	uchar8 d1 = src1[i1];
	uchar8 d2 = src1[i2];
	uchar8 d3 = src2[i1];
	uchar8 d4 = src2[i2];

	avr1.s0 = (d1.s0 + d1.s1 + d2.s0 + d2.s1)/4;
	avr1.s1 = (d1.s2 + d1.s3 + d2.s2 + d2.s3)/4;
	avr1.s2 = (d1.s4 + d1.s5 + d2.s4 + d2.s5)/4;
	avr1.s3 = (d1.s6 + d1.s7 + d2.s6 + d2.s7)/4;
                                      
	avr2.s0 = (d3.s0 + d3.s1 + d4.s0 + d4.s1)/4;
	avr2.s1 = (d3.s2 + d3.s3 + d4.s2 + d4.s3)/4;
	avr2.s2 = (d3.s4 + d3.s5 + d4.s4 + d4.s5)/4;
	avr2.s3 = (d3.s6 + d3.s7 + d4.s6 + d4.s7)/4;
                                      
	dst1[y*SIZEX/8+x] = avr1;
	dst2[y*SIZEX/8+x] = avr2;
}

__kernel void subsample1(
					__global const uchar8 *src1,
					__global const uchar8 *src2,
					__global uchar4 *dst1,
					__global uchar4 *dst2)
{
	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t i1 = y*2*SIZEX/8 + x;	// first line
	size_t i2 = i1 + SIZEX/8;		// secons line

	ushort4 avr1;
	ushort4 avr2;

	uchar8 d1 = src1[i1];
	avr1.s0 = (d1.s0 + d1.s1);
	avr1.s1 = (d1.s2 + d1.s3);
	avr1.s2 = (d1.s4 + d1.s5);
	avr1.s3 = (d1.s6 + d1.s7);

	uchar8 d2 = src1[i2];
	avr1.s0 += (d2.s0 + d2.s1);
	avr1.s1 += (d2.s2 + d2.s3);
	avr1.s2 += (d2.s4 + d2.s5);
	avr1.s3 += (d2.s6 + d2.s7);
                                      
	uchar8 d3 = src2[i1];
	avr2.s0 = (d3.s0 + d3.s1);
	avr2.s1 = (d3.s2 + d3.s3);
	avr2.s2 = (d3.s4 + d3.s5);
	avr2.s3 = (d3.s6 + d3.s7);
                                      
	uchar8 d4 = src2[i2];
	avr2.s0 += (d4.s0 + d4.s1);
	avr2.s1 += (d4.s2 + d4.s3);
	avr2.s2 += (d4.s4 + d4.s5);
	avr2.s3 += (d4.s6 + d4.s7);
              
	dst1[y*SIZEX/8+x] = convert_uchar4(avr1/(ushort4)(4));
	dst2[y*SIZEX/8+x] = convert_uchar4(avr2/(ushort4)(4));
}

__kernel void subsample0(
					__global const uchar8 *src,
					__global uchar4 *dst)
{
	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t i1 = SIZEY*SIZEX/8 + y*2*SIZEX/8 + x;	// first line
	size_t i2 = i1 + SIZEX/8;						// secons line

	ushort4 avr;

	uchar8 d1 = src[i1];
	avr.s0 = (d1.s0 + d1.s1);
	avr.s1 = (d1.s2 + d1.s3);
	avr.s2 = (d1.s4 + d1.s5);
	avr.s3 = (d1.s6 + d1.s7);

	uchar8 d2 = src[i2];
	avr.s0 += (d2.s0 + d2.s1);
	avr.s1 += (d2.s2 + d2.s3);
	avr.s2 += (d2.s4 + d2.s5);
	avr.s3 += (d2.s6 + d2.s7);
                                      
	dst[y*SIZEX/8+x] = convert_uchar4(avr/(ushort4)(4));
}

// 8x8 order
__kernel void subsample3(
					__global const uchar16 *src,
					__global uchar8 *dst)
{
	size_t i = get_global_id(0);
	size_t i1 = SIZEY*SIZEX/16 + i*2;	// first line
	size_t i2 = i1 + 1;					// secons line

	ushort8 avr;

	uchar16 d1 = src[i1];
	avr.s0 = (d1.s0 + d1.s1);
	avr.s1 = (d1.s2 + d1.s3);
	avr.s2 = (d1.s4 + d1.s5);
	avr.s3 = (d1.s6 + d1.s7);
	avr.s4 = (d1.s8 + d1.s9);
	avr.s5 = (d1.sA + d1.sB);
	avr.s6 = (d1.sC + d1.sD);
	avr.s7 = (d1.sE + d1.sF);

	uchar16 d2 = src[i2];
	avr.s0 += (d2.s0 + d2.s1);
	avr.s1 += (d2.s2 + d2.s3);
	avr.s2 += (d2.s4 + d2.s5);
	avr.s3 += (d2.s6 + d2.s7);
	avr.s4 += (d1.s8 + d1.s9);
	avr.s5 += (d1.sA + d1.sB);
	avr.s6 += (d1.sC + d1.sD);
	avr.s7 += (d1.sE + d1.sF);
                                      
	dst[i] = convert_uchar8(avr/(ushort8)(4));
}

__kernel void trans0(__global uchar8 *src)
{
	size_t g0 = get_global_id(0);
	size_t g1 = get_global_id(1);
	size_t l = get_local_id(1);
	size_t i = g1*SIZEX/8 + g0;

	__local uchar d[8][8];

	d[0][l] = src[i].s0;
	d[1][l] = src[i].s1;
	d[2][l] = src[i].s2;
	d[3][l] = src[i].s3;
	d[4][l] = src[i].s4;
	d[5][l] = src[i].s5;
	d[6][l] = src[i].s6;
	d[7][l] = src[i].s7;

	barrier(CLK_LOCAL_MEM_FENCE);

	src[i].s0 = d[l][0];
	src[i].s1 = d[l][1];
	src[i].s2 = d[l][2];
	src[i].s3 = d[l][3];
	src[i].s4 = d[l][4];
	src[i].s5 = d[l][5];
	src[i].s6 = d[l][6];
	src[i].s7 = d[l][7];
}                  

__kernel void dct3(__global uchar8 *pixels, __global short8 *data)
{
	const int // Ci = cos(i*PI/16)*(1<<14);
		C1 = 16070,
		C2 = 15137,
		C3 = 13623,
		C4 = 11586,
		C5 = 9103,
		C6 = 6270,
		C7 = 3197;
/*
	size_t g = get_global_id(0);
	size_t l = g % 8;
	size_t n = g / SIZEX;
	size_t i = n*(SIZEX-SIZEX/8) + g/8 + l*(SIZEX/8);
*/
	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t l = get_local_id(1);
	size_t n = y / 8;
	size_t i = y*get_global_size(0) + x;
	size_t g = n*get_global_size(0)*8 + x*8 + l;
	size_t m = get_local_id(0);

	__local int rows[8][8][9];

	int s07,s16,s25,s34,s0734,s1625;
	int d07,d16,d25,d34,d0734,d1625;

	short8 p = convert_short8(pixels[i]) - (short8)(128);

	s07 = p.s0 + p.s7;
	d07 = p.s0 - p.s7;
	s16 = p.s1 + p.s6;
	d16 = p.s1 - p.s6;
	s25 = p.s2 + p.s5;
	d25 = p.s2 - p.s5;
	s34 = p.s3 + p.s4;
	d34 = p.s3 - p.s4;

	rows[m][1][l] = (C1*d07 + C3*d16 + C5*d25 + C7*d34) >> 14;
	rows[m][3][l] = (C3*d07 - C7*d16 - C1*d25 - C5*d34) >> 14;
	rows[m][5][l] = (C5*d07 - C1*d16 + C7*d25 + C3*d34) >> 14;
	rows[m][7][l] = (C7*d07 - C5*d16 + C3*d25 - C1*d34) >> 14;
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	rows[m][0][l] = (C4*(s0734 + s1625)) >> 14;
	rows[m][4][l] = (C4*(s0734 - s1625)) >> 14;
          
	rows[m][2][l] = (C2*d0734 + C6*d1625) >> 14;
	rows[m][6][l] = (C6*d0734 - C2*d1625) >> 14;

	barrier(CLK_LOCAL_MEM_FENCE);

	s07 = rows[m][l][0] + rows[m][l][7];
	d07 = rows[m][l][0] - rows[m][l][7];
	s16 = rows[m][l][1] + rows[m][l][6];
	d16 = rows[m][l][1] - rows[m][l][6];
	s25 = rows[m][l][2] + rows[m][l][5];
	d25 = rows[m][l][2] - rows[m][l][5];
	s34 = rows[m][l][3] + rows[m][l][4];
	d34 = rows[m][l][3] - rows[m][l][4];
                             
	data[g].s1 = (C1*d07 + C3*d16 + C5*d25 + C7*d34) >> 16;
	data[g].s3 = (C3*d07 - C7*d16 - C1*d25 - C5*d34) >> 16;
	data[g].s5 = (C5*d07 - C1*d16 + C7*d25 + C3*d34) >> 16;
	data[g].s7 = (C7*d07 - C5*d16 + C3*d25 - C1*d34) >> 16;
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	data[g].s0 = (C4*(s0734 + s1625)) >> 16;
	data[g].s4 = (C4*(s0734 - s1625)) >> 16;
          
	data[g].s2 = (C2*d0734 + C6*d1625) >> 16;
	data[g].s6 = (C6*d0734 - C2*d1625) >> 16;
}

__kernel void dct3f(__global uchar8 *pixels, __global short8 *data)
{
	const float // Ci = cos(i*PI/16);
		C1 = 0.98078528040323044912618223613424,
		C2 = 0.92387953251128675612818318939679,
		C3 = 0.83146961230254523707878837761791,
		C4 = 0.70710678118654752440084436210485,
		C5 = 0.55557023301960222474283081394853,
		C6 = 0.3826834323650897717284599840304,
		C7 = 0.19509032201612826784828486847702;

	size_t x = get_global_id(0);
	size_t y = get_global_id(1);
	size_t l = get_local_id(1);
	size_t n = y / 8;
	size_t i = y*get_global_size(0) + x;
	size_t g = n*get_global_size(0)*8 + x*8 + l;
	size_t m = get_local_id(0);

	__local float rows[8][8][9];

	float s07,s16,s25,s34,s0734,s1625;
	float d07,d16,d25,d34,d0734,d1625;

	float8 p = convert_float8(pixels[i]) - 128.0f;

	s07 = p.s0 + p.s7;
	d07 = p.s0 - p.s7;
	s16 = p.s1 + p.s6;
	d16 = p.s1 - p.s6;
	s25 = p.s2 + p.s5;
	d25 = p.s2 - p.s5;
	s34 = p.s3 + p.s4;
	d34 = p.s3 - p.s4;

	rows[m][1][l] = (C1*d07 + C3*d16 + C5*d25 + C7*d34);
	rows[m][3][l] = (C3*d07 - C7*d16 - C1*d25 - C5*d34);
	rows[m][5][l] = (C5*d07 - C1*d16 + C7*d25 + C3*d34);
	rows[m][7][l] = (C7*d07 - C5*d16 + C3*d25 - C1*d34);
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	rows[m][0][l] = (C4*(s0734 + s1625));
	rows[m][4][l] = (C4*(s0734 - s1625));
          
	rows[m][2][l] = (C2*d0734 + C6*d1625);
	rows[m][6][l] = (C6*d0734 - C2*d1625);

	barrier(CLK_LOCAL_MEM_FENCE);

	s07 = rows[m][l][0] + rows[m][l][7];
	d07 = rows[m][l][0] - rows[m][l][7];
	s16 = rows[m][l][1] + rows[m][l][6];
	d16 = rows[m][l][1] - rows[m][l][6];
	s25 = rows[m][l][2] + rows[m][l][5];
	d25 = rows[m][l][2] - rows[m][l][5];
	s34 = rows[m][l][3] + rows[m][l][4];
	d34 = rows[m][l][3] - rows[m][l][4];
                             
	data[g].s1 = (C1*d07 + C3*d16 + C5*d25 + C7*d34)/4;
	data[g].s3 = (C3*d07 - C7*d16 - C1*d25 - C5*d34)/4;
	data[g].s5 = (C5*d07 - C1*d16 + C7*d25 + C3*d34)/4;
	data[g].s7 = (C7*d07 - C5*d16 + C3*d25 - C1*d34)/4;
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	data[g].s0 = (C4*(s0734 + s1625))/4;
	data[g].s4 = (C4*(s0734 - s1625))/4;
          
	data[g].s2 = (C2*d0734 + C6*d1625)/4;
	data[g].s6 = (C6*d0734 - C2*d1625)/4;
}

__kernel void dct3N(__global uchar8 *pixels, __global short8 *data)
{
	const int // Ci = cos(i*PI/16)*(1<<14);
		C1 = 16070,
		C2 = 15137,
		C3 = 13623,
		C4 = 11586,
		C5 = 9103,
		C6 = 6270,
		C7 = 3197;

	size_t i = get_global_id(0);
	size_t m = get_local_id(0)/8;
	size_t l = get_local_id(0)%8;

	__local int rows[8][8][9];

	int s07,s16,s25,s34,s0734,s1625;
	int d07,d16,d25,d34,d0734,d1625;

	short8 p = convert_short8(pixels[i]) - (short8)(128);

	s07 = p.s0 + p.s7;
	d07 = p.s0 - p.s7;
	s16 = p.s1 + p.s6;
	d16 = p.s1 - p.s6;
	s25 = p.s2 + p.s5;
	d25 = p.s2 - p.s5;
	s34 = p.s3 + p.s4;
	d34 = p.s3 - p.s4;

	rows[m][1][l] = (C1*d07 + C3*d16 + C5*d25 + C7*d34) >> 14;
	rows[m][3][l] = (C3*d07 - C7*d16 - C1*d25 - C5*d34) >> 14;
	rows[m][5][l] = (C5*d07 - C1*d16 + C7*d25 + C3*d34) >> 14;
	rows[m][7][l] = (C7*d07 - C5*d16 + C3*d25 - C1*d34) >> 14;
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	rows[m][0][l] = (C4*(s0734 + s1625)) >> 14;
	rows[m][4][l] = (C4*(s0734 - s1625)) >> 14;
          
	rows[m][2][l] = (C2*d0734 + C6*d1625) >> 14;
	rows[m][6][l] = (C6*d0734 - C2*d1625) >> 14;

	barrier(CLK_LOCAL_MEM_FENCE);

	s07 = rows[m][l][0] + rows[m][l][7];
	d07 = rows[m][l][0] - rows[m][l][7];
	s16 = rows[m][l][1] + rows[m][l][6];
	d16 = rows[m][l][1] - rows[m][l][6];
	s25 = rows[m][l][2] + rows[m][l][5];
	d25 = rows[m][l][2] - rows[m][l][5];
	s34 = rows[m][l][3] + rows[m][l][4];
	d34 = rows[m][l][3] - rows[m][l][4];
                             
	data[i].s1 = (C1*d07 + C3*d16 + C5*d25 + C7*d34) >> 16;
	data[i].s3 = (C3*d07 - C7*d16 - C1*d25 - C5*d34) >> 16;
	data[i].s5 = (C5*d07 - C1*d16 + C7*d25 + C3*d34) >> 16;
	data[i].s7 = (C7*d07 - C5*d16 + C3*d25 - C1*d34) >> 16;
          
	s0734 = s07 + s34;
	d0734 = s07 - s34;
	s1625 = s16 + s25;
	d1625 = s16 - s25;

	data[i].s0 = (C4*(s0734 + s1625)) >> 16;
	data[i].s4 = (C4*(s0734 - s1625)) >> 16;
          
	data[i].s2 = (C2*d0734 + C6*d1625) >> 16;
	data[i].s6 = (C6*d0734 - C2*d1625) >> 16;
}
